//------------------------头文件部分------------------------
#include <bits/stdc++.h>         // 包含C++标准库（实际开发中不推荐使用）
#include <cuda.h>                // CUDA主头文件
#include "cuda_runtime.h"        // CUDA运行时API
#include "device_launch_parameters.h" // 设备启动参数
#include <time.h>                // 时间相关函数
#include <sys/time.h>            // 高精度时间测量

//----------------------宏定义部分--------------------------
#define THREAD_PER_BLOCK 256     // 定义每个块的线程数

//---------------------CUDA核函数部分-----------------------
__global__ void reduce0(float *d_in, float *d_out, unsigned int N) {
    
    __shared__ float sdata[THREAD_PER_BLOCK];  // 声明共享内存数组

    // 线程索引计算
    unsigned int tid = threadIdx.x;            // 线程在块内的ID
    



    float* inputbegin=d_in+blockIdx.x*blockDim.x ;

    

    // 归约计算阶段（低效实现）
    for(unsigned int s=1; s < blockDim.x; s *= 2) {  // 步长倍增循环
        if (tid % (2*s) == 0) {  // 仅选择特定线程进行计算（导致线程束分化）
            //sdata[tid] += sdata[tid + s]; // 相邻元素相加
            inputbegin[tid]+=inputbegin[tid+s];
        }
        __syncthreads();         // 每次迭代后同步
    }

    // 结果写回阶段
    if (tid == 0)                // 仅块内第一个线程执行
         

        d_out[blockIdx.x] = inputbegin[0]; // 将块结果写入全局内存
}


__global__ void reducedeepseek(float *d_in, float *d_out, unsigned int N) {
    
    __shared__ float sdata[THREAD_PER_BLOCK];  // 声明共享内存数组

    // 线程索引计算
    unsigned int tid = threadIdx.x;            // 线程在块内的ID
    unsigned int i = blockIdx.x*blockDim.x + threadIdx.x; // 全局内存索引



   

    // 数据加载阶段：每个线程加载一个元素到共享内存
    sdata[tid] = d_in[i];        // 潜在问题：当i>=N时会越界访问
        // 修复问题1和2：添加边界检查
    sdata[tid] = (i < N) ? d_in[i] : 0.0f;
    __syncthreads();             // 块内同步确保所有线程完成数据加载

    // 归约计算阶段（低效实现）
    for(unsigned int s=1; s < blockDim.x; s *= 2) {  // 步长倍增循环
        if (tid % (2*s) == 0) {  // 仅选择特定线程进行计算（导致线程束分化）
            sdata[tid] += sdata[tid + s]; // 相邻元素相加
           
        }
        __syncthreads();         // 每次迭代后同步
    }

    // 结果写回阶段
    if (tid == 0)                // 仅块内第一个线程执行
         d_out[blockIdx.x] = sdata[0]; // 将块结果写入全局内存

        
}

// reducedeepseek优势：

// 利用共享内存降低全局内存访问开销。

// 通过边界检查提升安全性。

// 更适合大规模数据归约。

// reduce0缺陷：

// 全局内存操作导致性能低下。

// 潜在越界风险和同步问题。

// 结论：reducedeepseek在性能、安全性上全面优于reduce0，但两者仍需进一步优化线程束分化和访存模式。

//-------------------结果验证函数部分----------------------
bool check(float *out, float *res, int n) {
    for(int i=0; i<n; i++) {
        // 绝对误差检查（建议改用相对误差检查）
        if(abs(out[i]-res[i])>0.005) // 0.005的阈值可能不够严谨
            return false;
    }
    return true;
}

//---------------------主函数部分--------------------------
int main() {
    // 问题规模定义
    const int N = 32*1024*1024;  // 32M个浮点数

    // 主机内存分配
    float *input = (float *)malloc(N*sizeof(float));    // 输入数据
    float *output = (float *)malloc((N/THREAD_PER_BLOCK)*sizeof(float)); // GPU输出
    float *result = (float *)malloc((N/THREAD_PER_BLOCK)*sizeof(float)); // CPU参考结果

    // 设备内存分配
    float *d_input, *d_output;
    cudaMalloc((void **)&d_input, N*sizeof(float));                 // 设备输入
    cudaMalloc((void **)&d_output, (N/THREAD_PER_BLOCK)*sizeof(float)); // 设备输出

    if (!d_input|| !d_output ) { // 检查分配是否成功
        printf("GPU内存分配失败\n");
        return -1;
    }

    // 数据初始化（未设置随机种子）
    for(int i=0; i<N; i++) {
        input[i] = 2.0*(float)drand48()-1.0;  // 生成[-1,1)范围内的随机数
    }

    // CPU参考计算（假设N是THREAD_PER_BLOCK的整数倍）
    int block_num = N/THREAD_PER_BLOCK;
    for(int i=0; i<block_num; i++) {
        float cur = 0;
        for(int j=0; j<THREAD_PER_BLOCK; j++) {  // 无越界保护
            cur += input[i*THREAD_PER_BLOCK + j];
        }
        result[i] = cur;
    }

    // 数据传输：主机->设备（缺少错误检查）
   
    // 所有cudaMemcpy调用后添加检查
    cudaError_t err =  cudaMemcpy(d_input, input, N*sizeof(float), cudaMemcpyHostToDevice);
    if (err != cudaSuccess) {
         printf("Memcpy failed: %s\n", cudaGetErrorString(err));
    }

    // 核函数配置
    dim3 Grid(N/THREAD_PER_BLOCK, 1);  // 一维网格（假设N可被整除）
    dim3 Block(THREAD_PER_BLOCK, 1);   // 一维块

    // 核函数调用（缺少错误检查）
    reduce0<<<Grid, Block>>>(d_input, d_output,N);

    // 数据传输：设备->主机（缺少错误检查）
    cudaError_t err2 = cudaMemcpy(output, d_output, block_num*sizeof(float), cudaMemcpyDeviceToHost);
    if (err2 != cudaSuccess) {
         printf("Memcpy failed: %s\n", cudaGetErrorString(err2));
    }

    // 结果验证
    if(check(output, result, block_num)) {
        printf("the ans is right\n");
    } else {
        printf("the ans is wrong\n");
        // 打印全部错误结果（大数据量时可能不适用）
        for(int i=0; i<block_num; i++) {
            printf("%lf ", output[i]);
        }
        printf("\n");
    }

    // 资源释放
    free(input);
    free(output);
    free(result);
    cudaFree(d_input);
    cudaFree(d_output);
    
}

// 核函数问题：

// 归约算法使用tid % (2*s)条件，导致严重的线程束分化

// 共享内存访问模式不是合并访问，影响内存效率

// 没有处理N不是线程块整数倍的情况

// 内存管理问题：

// 缺少cudaMalloc/malloc的返回值检查

// 主机端分配的内存input/output/result没有释放

// 设备到主机的数据传输没有错误检查

// 功能限制：

// 验证函数使用绝对误差检查，对大规模数据不够可靠

// 没有初始化随机数种子（srand48缺失）

// 结果打印在大数据量时可能产生过多输出

// 潜在风险：

// 当N % THREAD_PER_BLOCK != 0时会出现内存越界

// 没有使用cudaDeviceSynchronize()确保核函数完成
